




__kernel void granularity_test_kernel(__global int* src, __global int* dst) {


	int id = get_global_id(0);
	int data = *(src + id);
	
	if (data % 2) {
		data = data*data; //the short branch
	} else {
		data = exp( log( exp( log((float)data) ) ) ); //the long branch
		data = exp( log( exp( log((float)data) ) ) );
		data = exp( log( exp( log((float)data) ) ) );
		data = exp( log( exp( log((float)data) ) ) );
		data = exp( log( exp( log((float)data) ) ) );
		data = exp( log( exp( log((float)data) ) ) );
		data = exp( log( exp( log((float)data) ) ) );
		data = exp( log( exp( log((float)data) ) ) );
		data = exp( log( exp( log((float)data) ) ) );
		data = exp( log( exp( log((float)data) ) ) );
		data = exp( log( exp( log((float)data) ) ) );
		data = exp( log( exp( log((float)data) ) ) );
		data = exp( log( exp( log((float)data) ) ) );
		data = exp( log( exp( log((float)data) ) ) );
		data = exp( log( exp( log((float)data) ) ) );
		data = exp( log( exp( log((float)data) ) ) );
		data = exp( log( exp( log((float)data) ) ) );
	}
	
	
	*(dst + id) = data/data + id;

}



//basic test kernel
__kernel void test_kernel(__global float* input, __global float* output) {

	int id = get_global_id(0);
	
	*(output + id) = id*id;
};






#define FS_BLOCK_X 64
#define FS_PIX_PER_THREAD 4
#define FS_MBLOCK_PER_CUDABLOCK ((FS_BLOCK_X * FS_PIX_PER_THREAD)/ 16)

#define FS_CHECK_POINTS 1024

__kernel void test_fullsearch_lowerbound(__global unsigned char* frame, unsigned int frameStride, __global unsigned char* reference, unsigned int referenceStride, __global int* output, unsigned int outputStride) {

    const int mb_x = FS_MBLOCK_PER_CUDABLOCK*get_group_id(0) + get_local_id(0)/16;
    const int mb_y = get_group_id(1);
    
    __local unsigned char shared_block[ 256 * FS_MBLOCK_PER_CUDABLOCK ];
    __local unsigned char shared_reference[ 256 * FS_MBLOCK_PER_CUDABLOCK ];
    __local int shared_sad[ FS_BLOCK_X ];
    
    
   // #pragma unroll
    for (int i = 0; i < 16*FS_PIX_PER_THREAD; i++) {
    	int x_shift = (i % FS_PIX_PER_THREAD)*FS_BLOCK_X;
    	int y_shift = i / FS_PIX_PER_THREAD;
    	unsigned char color = *(frame + 16*mb_x + get_local_id(0) + x_shift + (16*mb_y + y_shift)*frameStride);
	    shared_block[get_local_id(0)*FS_PIX_PER_THREAD + x_shift + FS_BLOCK_X*FS_PIX_PER_THREAD*y_shift] = color;
    }
    
  //  #pragma unroll

    for (int i = 0; i < 16*FS_PIX_PER_THREAD; i++) {
  		int x_shift = (i % FS_PIX_PER_THREAD)*FS_BLOCK_X;
    	int y_shift = i / FS_PIX_PER_THREAD;
    	unsigned char color = *(frame + 16*mb_x + get_local_id(0) + x_shift + (16*mb_y + y_shift)*frameStride);
	    shared_reference[get_local_id(0)*FS_PIX_PER_THREAD + x_shift + FS_BLOCK_X*FS_PIX_PER_THREAD*y_shift] = color ;
    }
    
/*
	//a sad that may cause bank conflicts in shared memory
    int sad_acum = 0;
    for(int k = 0; k < 1024; k++) {
    //	#pragma unroll
    	for (int i = 0; i < 16*FS_PIX_PER_THREAD; i++) {
    		int x_shift = i % FS_PIX_PER_THREAD;
    		int y_shift = i / FS_PIX_PER_THREAD;
    		int diff = abs(shared_reference[get_local_id(0)*FS_PIX_PER_THREAD + x_shift + FS_BLOCK_X*FS_PIX_PER_THREAD*y_shift] - shared_block[get_local_id(0)*FS_PIX_PER_THREAD + x_shift + FS_BLOCK_X*FS_PIX_PER_THREAD*y_shift] );
    		sad_acum += diff;
    	}
    
    	shared_sad[ get_local_id(0) ] = sad_acum;
    	if (get_local_id(0) % FS_PIX_PER_THREAD == 0) {
    	//	#pragma unroll
    		for (int i = 1; i < FS_PIX_PER_THREAD; i++) {
	    	sad_acum += shared_sad[ get_local_id(0) + i];
    		}
    	}
    	barrier(CLK_LOCAL_MEM_FENCE | CLK_GLOBAL_MEM_FENCE);
    }
   */
   
   	//a INVALID/INCORRECT/WRONG sad that simulates no bank conflict
	int sad_acum = 0;
    for(int k = 0; k < FS_CHECK_POINTS; k++) {
    	#pragma unroll
    	for (int i = 0; i < 16*FS_PIX_PER_THREAD; i++) {
    	//	int x_shift = i % FS_PIX_PER_THREAD;
    	//	int y_shift = i / FS_PIX_PER_THREAD;
    		int diff = abs(shared_reference[get_local_id(0)*4] - shared_block[get_local_id(0)*4] ); //multiple of int
    		sad_acum += diff;
    	}
    
    	shared_sad[ get_local_id(0) ] = sad_acum;
    	if (get_local_id(0) % FS_PIX_PER_THREAD == 0) {
    		#pragma unroll
    		for (int i = 1; i < FS_PIX_PER_THREAD; i++) {
	    	sad_acum += shared_sad[ get_local_id(0) + i];
    		}
    	}
    	barrier(CLK_LOCAL_MEM_FENCE | CLK_GLOBAL_MEM_FENCE);
    }
    
    
    if (get_local_id(0) % FS_PIX_PER_THREAD == 0) {
	    output[mb_x + mb_y*outputStride ] = sad_acum;
    }
}


